#include "utility.h"

void copyparamtoGPU( FlipConstant hparam );

__global__ void cptdivergence( farray outdiv, farray ux, farray uy, farray uz, charray mark );
__global__ void advectparticle( float3 *pos, float3 *vel, int pnum, farray ux, farray uy, farray uz, float dt, char *parflag, VELOCITYMODEL velmode );
__global__ void advectparticle_RK2( float3 *ppos, float3 *pvel, int pnum, farray ux, farray uy, farray uz, float dt, 
								   char *parflag, VELOCITYMODEL velmode );
__global__ void mapvelp2g_slow( float3 *pos, float3 *vel, int pnum, farray ux, farray uy, farray uz );
__global__ void mapvelp2g_hash( float3 *pos, float3 *vel,  float *mass, char *parflag, int pnum, farray ux, farray uy, farray uz, uint* gridstart, uint *gridend );
__global__ void mapvelg2p_flip( float3 *pos, float3 *vel, char* parflag, int pnum, farray ux, farray uy, farray uz );
__global__ void subGradPress( farray p, farray ux, farray uy, farray uz );
__global__ void JacobiIter( farray outp, farray p, farray b, charray mark );
__global__ void markair( charray mark );
__global__ void markfluid( charray mark, float3 *pos, char *parflag, int pnum );
__global__ void markfluid_dense( charray mark, float *parmass, char *parflag, int pnum, uint *gridstart, uint *gridend, int fluidParCntPerGridThres );
__global__ void markBoundaryCell( charray mark );
__global__ void setgridcolor_k( float* color, ECOLORMODE mode, farray p, farray ux, farray uy, farray uz, farray div, farray phi, charray mark, farray dense );
__global__ void initphi( farray phi, charray mark );
__global__ void sweepphi( farray phi, charray mark );
__global__ void sweepu(farray outux, farray outuy, farray outuz, farray ux, farray uy, farray uz, farray phi, charray mark );
__global__ void setWaterBoundaryU_k( farray ux, farray uy, farray uz, charray mark );
__global__ void setSmokeBoundaryU_k( farray ux, farray uy, farray uz, charray mark );
__global__ void computeDeltaU( farray ux, farray uy, farray uz, farray uxold, farray uyold, farray uzold );
__global__ void addexternalforce_k( float3 *vel, int pnum, float dt );
__global__ void setPressBoundary( farray press );
__global__ void correctparticlepos( float3* outpos, float3* ppos, float *pmass,  char* parflag, int pnum, uint* gridstart, uint *gridend, float correctionspring, float correctionradius );
__global__ void copyParticle2GL_radius_k( float3* ppos, float *pmass, char *pflag, int pnum, float *renderpos, float *rendercolor, float minmass);
__global__ void copyParticle2GL_vel_k( float3* ppos, float3 *pvel, float *pmass, char *pflag, int pnum, float *renderpos, float *rendercolor );

//for smoke
__global__ void advectux( farray outux, farray ux, farray uy, farray uz, float velocitydissipation, float3 wind );
__global__ void advectuy( farray outuy, farray ux, farray uy, farray uz, float velocitydissipation, float3 wind );
__global__ void advectuz( farray outuz, farray ux, farray uy, farray uz, float velocitydissipation, float3 wind );
__global__ void advectscaler( farray outscalar, farray scalar, farray ux, farray uy, farray uz, float densedissipation, float3 wind );
__global__ void markforsmoke( charray mark );
__global__ void setsmokevel( farray uz, farray dense );
__global__ void setsmokedense( farray dense );
void writedens2surface( cudaArray* cudaarray, int blocknum, int threadnum, farray dense );
__global__ void diffuse_dense( farray outp, farray inp, charray mark, float alpha, float beta );
__global__ void diffuse_velocity( farray outv, farray inv, float alpha, float beta );

//pcg: preconditioned conjugate gradient method
__global__ void arrayproduct_k( float* out, float* x, float *y, int n );
__global__ void computeAx( farray ans, charray mark, farray x, int n );
__global__ void pcg_op( charray A, farray ans, farray x, farray y, float a, int n );
__global__ void applyprecondition_pcg( farray P, charray mark, farray ans, farray input, int n );

//for sort the particles.
__global__ void calcHashD(uint*   gridParticleHash, uint*   gridParticleIndex, float3* pos, uint numParticles);
__global__ void reorderDataAndFindCellStartD(uint*   cellStart,        // output: cell start index
											 uint*   cellEnd,          // output: cell end index
											 float3* sortedPos,        // output: sorted positions
											 float3* sortedVel,        // output: sorted velocities
											 float3* sortedVelLag,        // output: sorted velocities
											 char* sortedflag,
											 float* sortedmass,
											 float* sortedslipparam,
											 uint *  gridParticleHash, // input: sorted grid hashes
											 uint *  gridParticleIndex,// input: sorted particle indices
											 float3* oldPos,           // input: sorted position array
											 float3* oldVel,           // input: sorted velocity array
											 float3* oldVelLag,           // input: sorted velocity array
											 char* oldflag,
											 float* oldmass,
											 float* oldslipparam,
											 uint    numParticles);

//Marching Cube.
void allocateTextures( uint **d_edgeTable, uint **d_triTable,  uint **d_numVertsTable );
// compact voxel array
__global__ void genWaterDensfield( farray outdens, float3 *pos, char *parflag, uint *gridstart, uint  *gridend, float fMCDensity );
__global__ void genWaterDensfield2( farray outdens, float3 *pos, char *parflag, uint *gridstart, uint  *gridend, float fMCDensity );
__global__ void classifyVoxel(uint* voxelVerts, uint *voxelOccupied, farray volume, float isoValue);
__global__ void compactVoxels(uint *compactedVoxelArray, uint *voxelOccupied, uint *voxelOccupiedScan, uint numVoxels);
__global__ void generateTriangles2(float3 *pos, float3 *norm, uint *compactedVoxelArray, uint *numVertsScanned, farray volume,
				   float isoValue, uint activeVoxels, uint maxVerts);
__global__ void markActiveEdge_MC( uint *outmark, uint *compactedVoxelArray, farray volume, float isoValue, uint activeVoxels );
__global__ void generateTriangles_indices( float3 *pTriVertex, uint *pTriIndices, uint *compactedVoxelArray, farray volume,
						  float isoValue, uint activeVoxels, uint maxVerts, uint *MCEdgeIdxMapped, uint *vertexnumscan );
__global__ void calnormal_k( float3 *ppos, float3 *pnor, int pnum, uint *indices, int indicesnum );
__global__ void normalizeTriangleNor_k( float3 *pnor, int pnum );

//smoothing 3d triangle mesh
__global__ void smooth_computedisplacement( float3 *displacement, int *weight, float3 *ppos, uint *indices, int trianglenum );
__global__ void smooth_addDisplacement( float3 *displacement, int *weight, float3 *ppos, int vertexnum, float param );

//solid coupling
__global__ void markSolid_sphere( float3 spherepos, float sphereradius, charray mark );
